Fix uninitialized CUB reduction identity corrupting fill_holes vertices - #41

Summary

CuMesh::fill_holes produces corrupted vertices whenever the CUB reduction identity happens to land on non-zero host stack memory. In practice the first fill_holes call in a process is usually fine and every subsequent call is corrupted, which makes this look like non-determinism but is actually deterministic per-call-index state.

This is a cross-platform correctness bug, not a Windows-specific one — the identity is constructed on the host, so it depends on host stack residue, not on the compiler or OS.

Root cause

fill_holes computes each hole-cap vertex with a segmented reduction over Vec3f:

// src/clean_up.cu
cub::DeviceSegmentedReduce::Sum(..., (Vec3f*)sums, num_holes, offsets, offsets + 1, stream);
inplace_div_kernel<<<...>>>(sums, counts, num_holes);

CUB materializes the reduction identity on the host as InitT{}. Vec3f is declared with a user-provided default constructor:

// src/dtypes.cuh
struct __align__(16) Vec3f {
    float x, y, z;
    __device__ __forceinline__ Vec3f();   // user-provided => Vec3f{} is NOT zero-initialized
    ...
};

Two things combine:

1. Because the default constructor is user-provided, Vec3f{} is value-initialization that calls the constructor rather than zero-filling.
2. The constructor is marked __device__ only, so it is not callable from host code. CUB's host-side InitT{} therefore yields whatever bytes were on the host stack.

That garbage G is then folded into every segment's sum and divided by the segment size, so each bad vertex gets (G + Σ midpoints) / n.

Evidence

Dumping the corrupted vertex buffer showed all bad Z values were exactly G / n for small integers n, with identical mantissas at different exponents:

0x4F3D360B = 3.17443550e+09   (G/3)
0x4EBD360B = 1.58721775e+09   (G/6)
0x4E3D360B = 7.93608875e+08   (G/12)
...
G ≈ 9,523,306,496

A single constant divided by the segment size can only come from a non-zero reduction identity. Only the Z component was affected here — the stack bytes backing x/y happened to be zero in this process, which is exactly the "sometimes it works" signature.

The bad vertices were also the contiguous tail of the buffer (the newly appended hole-cap vertices), and face indices were all in range — consistent with fill_holes appending bad vertices rather than any earlier stage emitting them.

Fix

Make Vec3f's constructors __host__ __device__ so the host-side identity is actually constructed (and zeroed) instead of reading uninitialized stack memory. Six lines, no behavior change on the device side.

atlas.cu already uses the safe pattern for its reduction (float3 with an explicit make_float3(0, 0, 0) identity); this brings fill_holes in line with that.

QEM in the same header has the same __device__-only user-provided constructor pattern, but it is not used in any CUB reduction today, so it is intentionally left unchanged to keep this fix surgical.

Validation

Repro harness: N image-to-3D generations in a single process (TRELLIS.2), measuring max mesh edge length, long-edge count, and rendered mask coverage.

0/5 clean before, 5/5 clean after. Renders confirm it visually: the corrupted runs smeared the model into a cylindrical column of sliver triangles; after the fix every run reconstructs correctly.

Built and tested on Windows 11, MSVC 19.44, CUDA 13.0, PyTorch 2.13, RTX 5090 (sm_120).